// #define DEBUG_ENABLE_KUNLUN // 打印标量和矢量
#include "kernel_operator.h"
// #include "utils/copy_kl.h"
// #include "utils/debug_kl.h"

using namespace AscendC;
//const int32_t NN = 2;
class GroupNormKernelfloat
{
    using TYPE_X = float;
public:
    __aicore__ inline GroupNormKernelfloat() {} 

    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, int32_t num_groups, int32_t num_channels, int32_t channels_per_group,
                                float eps, bool affine, int32_t N, int32_t C, int32_t H, int32_t W)
    {
        this->num_groups = num_groups;
        this->num_channels = num_channels;
        this->channels_per_group = channels_per_group;
        this->eps = eps;
        this->affine = affine;
        this->N = N;
        this->C = C;
        this->H = H;
        this->W = W;
        xGm.SetGlobalBuffer((__gm__ TYPE_X *)x, N * C * H * W);
        yGm.SetGlobalBuffer((__gm__ TYPE_X *)y, N * C * H * W);
        int32_t NN = 8;
        if (TILING_KEY_IS(1)||TILING_KEY_IS(2))
        {
            pipe.InitBuffer(tmpbuf1, this->channels_per_group * H * sizeof(TYPE_X));
            pipe.InitBuffer(tmpbuf2, this->channels_per_group * H * sizeof(TYPE_X));
            pipe.InitBuffer(tmpbuf3, this->channels_per_group * H * sizeof(TYPE_X));
            pipe.InitBuffer(tmpbuf4, this->channels_per_group * H * sizeof(TYPE_X));
            pipe.InitBuffer(tmpbuf5, this->channels_per_group * H * sizeof(TYPE_X));
        }
        if (TILING_KEY_IS(3))
        {
            pipe.InitBuffer(tmpbuf1, NN * W * sizeof(TYPE_X));
            pipe.InitBuffer(tmpbuf2, NN * W * sizeof(TYPE_X));
            pipe.InitBuffer(tmpbuf3, NN * W * sizeof(TYPE_X));
            pipe.InitBuffer(tmpbuf4, NN * W * sizeof(TYPE_X));
            pipe.InitBuffer(tmpbuf5, NN * W * sizeof(TYPE_X));
        }
    }

    __aicore__ inline void Process1()
    {
        for (int32_t i = 0; i < N * num_groups; i++)
        {
            compute1(i);
        }
    }
    __aicore__ inline void Process2()
    {
        for (int32_t i = 0; i < N * num_groups; i++)
        {
            compute1(i);
        }
    }
    __aicore__ inline void Process3()
    {  
        for (int32_t i = 0; i < N * num_groups; i++)
        {   
            compute2(i);
        }
    }

private:
    __aicore__ inline void compute1(int32_t progress)
    {
        auto buf1 = tmpbuf1.Get<TYPE_X>();//x
        auto buf2 = tmpbuf2.Get<TYPE_X>();//ex
        auto buf3 = tmpbuf3.Get<TYPE_X>(); //var
        auto buf4 = tmpbuf4.Get<TYPE_X>();//eps
        auto buf5 = tmpbuf5.Get<TYPE_X>();//eps
        int32_t len = channels_per_group * H;
        TYPE_X mean,var,sum2;
        eps = 1e-5;
        DataCopyGM2UB<TYPE_X>(buf1, xGm[len * progress], len);
        AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
        AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
        ReduceSum<TYPE_X>(buf2, buf1, buf5, len);
        buf3(0) = len;
        Div(buf3, buf2, buf3, 1);
        mean = buf3(0);
        Mul(buf2, buf1, buf1, len);
        ReduceSum<TYPE_X>(buf4, buf2, buf5, len);
        sum2 = buf4(0);
        var = (sum2 - len * mean * mean) / len;
        Duplicate<TYPE_X>(buf3, var, len);
        AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
        AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
        //x-ex
        Duplicate<TYPE_X>(buf2, mean, len);
        Sub(buf1, buf1, buf2, len);
        //var+eps
        Duplicate<TYPE_X>(buf4, eps, len);
        Add(buf3, buf3, buf4, len);
        Sqrt(buf4, buf3, len);
        Div(buf1, buf1, buf4, len);
        AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
        AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
        DataCopyUB2GM<TYPE_X>(yGm[len * progress], buf1, len);
        AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0);
        AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0);
    }
private:
    __aicore__ inline void compute2(int32_t progress)
    {
        auto buf1 = tmpbuf1.Get<TYPE_X>();//x
        auto buf2 = tmpbuf2.Get<TYPE_X>();//ex
        auto buf3 = tmpbuf3.Get<TYPE_X>(); //var
        auto buf4 = tmpbuf4.Get<TYPE_X>();//eps
        auto buf5 = tmpbuf5.Get<TYPE_X>();//
        float sumNN = 0;
        float sum2 = 0;
        int32_t NN = 8;
        int32_t T = channels_per_group * H * W;
        int32_t h_groups = (channels_per_group * H) / NN;
        int32_t len = NN * W;
        int32_t lentail = ((channels_per_group * H) % NN) * W;
        int32_t lenTemp;
        int32_t Times = 0;
        if(lentail == 0){
            Times = h_groups;
        }
        else{
            Times = h_groups + 1;
        }
        float mean,var;
        eps = 1e-5;
        buf3(0) = sumNN;
        buf4(0) = sum2;
        for(int32_t j = 0; j < Times ; j++)
        {
            if(j == h_groups){
                lenTemp = lentail;
            }else{
                lenTemp = len;
            }
            DataCopyGM2UB<TYPE_X>(buf1, xGm[T * progress + j * len], lenTemp);
            AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
            ReduceSum<TYPE_X>(buf2, buf1, buf5, lenTemp);
            Add(buf3, buf2, buf3, 1);
            AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0);
            //jiade
            Mul(buf2, buf1, buf1, lenTemp);
            ReduceSum<TYPE_X>(buf5, buf2, buf2, lenTemp);
            Add(buf4, buf5, buf4, 1);
        }
        sumNN = (float)buf3(0);
        sum2 = (float)buf4(0);
        mean = sumNN / T;
        var = (sum2 - T * mean * mean) / T;
        // buf5(0) = sumval;
        // for(int32_t j = 0; j < Times; j++)
        // {
        //     if(j == h_groups){
        //         lenTemp = lentail;
        //     }else{
        //         lenTemp = len;
        //     }
        //     DataCopyGM2UB<TYPE_X>(buf1, xGm[channels_per_group * H * W * progress + j * len], lenTemp);
        //     AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
        //     AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
        //     Duplicate<TYPE_X>(buf2, mean, lenTemp);
        //     Sub(buf3, buf1, buf2, lenTemp);
        //     Mul(buf4, buf3, buf3, lenTemp);
        //     ReduceSum<TYPE_X>(buf3, buf4, buf4, lenTemp);
        //     Add(buf5, buf3, buf5, 1);
        //     AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0);
        //     AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0);
        // }
        // sumval = (float)buf5(0);
        // var = sumval / (channels_per_group * H * W);
        for(int32_t j = 0; j < Times; j++)
        {
            if(j == h_groups){
                lenTemp = lentail;
            }else{
                lenTemp = len;
            }
            Duplicate<TYPE_X>(buf3, var, len);
            AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
            DataCopyGM2UB<TYPE_X>(buf1, xGm[T * progress + j * len], lenTemp);
            AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
            //x-ex
            Duplicate<TYPE_X>(buf2, mean, lenTemp);
            Sub(buf1, buf1, buf2, lenTemp);
            Duplicate<TYPE_X>(buf4, eps, lenTemp);
            Add(buf3, buf3, buf4, lenTemp);
            Sqrt(buf4, buf3, lenTemp);
            Div(buf1, buf1, buf4, lenTemp);
            AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
            DataCopyUB2GM<TYPE_X>(yGm[T * progress + j * len], buf1, lenTemp);
            AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0);
        }

    }
    template <typename T>
    __aicore__ inline void DataCopyGM2UB(const LocalTensor<T> &dstTensor, const GlobalTensor<T> &srcTensor, const uint32_t &calCount)
    {
        DataCopyExtParams copyParams{1, calCount * static_cast<uint32_t>(sizeof(T)), 0, 0, 0};
        DataCopyPadExtParams<T> padParams{false, 0, 0, 0};

        DataCopyPad(dstTensor, srcTensor, copyParams, padParams);
    }

    template <typename T>
    __aicore__ inline void DataCopyUB2GM(const GlobalTensor<T> &dstTensor, const LocalTensor<T> &srcTensor, const uint32_t &calCount)
    {
        DataCopyExtParams copyParams{1, calCount * static_cast<uint32_t>(sizeof(T)), 0, 0, 0};
        DataCopyPad(dstTensor, srcTensor, copyParams);
    }
private:
    GlobalTensor<TYPE_X> xGm;
    GlobalTensor<TYPE_X> yGm;
    TPipe pipe;
    TBuf<QuePosition::VECCALC> tmpbuf1, tmpbuf2, tmpbuf3, tmpbuf4, tmpbuf5;

    int32_t N, C, H, W, DT;
    int32_t num_groups, num_channels, channels_per_group;
    float eps;
    bool affine;
};

class GroupNormKernelbf16
{
public:
    __aicore__ inline GroupNormKernelbf16() {} 

    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, int32_t num_groups, int32_t num_channels, int32_t channels_per_group,
                                float eps, bool affine, int32_t N, int32_t C, int32_t H, int32_t W)
    {
        this->num_groups = num_groups;
        this->num_channels = num_channels;
        this->channels_per_group = channels_per_group;
        this->eps = eps;
        this->affine = affine;
        this->N = N;
        this->C = C;
        this->H = H;
        this->W = W;
        xGm.SetGlobalBuffer((__gm__ bfloat16_t *)x, N * C * H * W);
        yGm.SetGlobalBuffer((__gm__ bfloat16_t *)y, N * C * H * W);
        int32_t NN = H;
        if (TILING_KEY_IS(4)||TILING_KEY_IS(5))
        {
            pipe.InitBuffer(tmpbuf1, this->channels_per_group * H * sizeof(float));
            pipe.InitBuffer(tmpbuf1bf16, this->channels_per_group * H * sizeof(bfloat16_t));
            pipe.InitBuffer(tmpbuf2, this->channels_per_group * H * sizeof(float));
            pipe.InitBuffer(tmpbuf3, this->channels_per_group * H * sizeof(float));
            pipe.InitBuffer(tmpbuf4, this->channels_per_group * H * sizeof(float));
            pipe.InitBuffer(tmpbuf5, this->channels_per_group * H * sizeof(float));
        }
        if (TILING_KEY_IS(6))
        {
            pipe.InitBuffer(tmpbuf1, H * W * sizeof(float));
            pipe.InitBuffer(tmpbuf1bf16, H * W * sizeof(bfloat16_t));
            pipe.InitBuffer(tmpbuf2, H * W * sizeof(float));
            pipe.InitBuffer(tmpbuf3, H * W * sizeof(float));
            pipe.InitBuffer(tmpbuf4, H * W * sizeof(float));
            pipe.InitBuffer(tmpbuf5, H * W * sizeof(float));
        }
    }

    __aicore__ inline void Process4()
    {
        for (int32_t i = 0; i < N * num_groups; i++)
        {
            compute3(i);
        }
    }
    __aicore__ inline void Process5()
    {
        for (int32_t i = 0; i < N * num_groups; i++)
        {
            compute3(i);
        }
    }
    __aicore__ inline void Process6()
    {  
        for (int32_t i = 0; i < N * num_groups; i++)
        {   
            compute4(i);
        }
    }

private:
    __aicore__ inline void compute3(int32_t progress)
    {
        auto buf1 = tmpbuf1.Get<float>();//x
        auto buf1bf16 = tmpbuf1bf16.Get<bfloat16_t>();
        auto buf2 = tmpbuf2.Get<float>();//ex
        auto buf3 = tmpbuf3.Get<float>(); //var
        auto buf4 = tmpbuf4.Get<float>();//eps
        auto buf5 = tmpbuf5.Get<float>();//eps
        int32_t len = channels_per_group * H;
        float mean,var,sum2;
        eps = 1e-5;
        DataCopyGM2UB<bfloat16_t>(buf1bf16, xGm[len * progress], len);
        AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
        AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
        Cast(buf1, buf1bf16, AscendC::RoundMode::CAST_NONE, len);
        ReduceSum<float>(buf2, buf1, buf5, len);
        buf3(0) = len;
        Div(buf3, buf2, buf3, 1);
        mean = buf3(0);

        Mul(buf2, buf1, buf1, len);
        ReduceSum<float>(buf4, buf2, buf5, len);
        sum2 = buf4(0);
        var = (sum2 - len * mean * mean) / len;
        Duplicate<float>(buf3, var, len);
        AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
        AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
        //x-ex
        Duplicate<float>(buf2, mean, len);
        Sub(buf1, buf1, buf2, len);
        //var+eps
        Duplicate<float>(buf4, eps, len);
        Add(buf3, buf3, buf4, len);
        Sqrt(buf4, buf3, len);
        Div(buf1, buf1, buf4, len);
        Cast(buf1bf16, buf1, AscendC::RoundMode::CAST_RINT, len);
        AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
        AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
        DataCopyUB2GM<bfloat16_t>(yGm[len * progress], buf1bf16, len);
        AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0);
        AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0);
    }
private:
    __aicore__ inline void compute4(int32_t progress)
    {
        auto buf1 = tmpbuf1.Get<float>();//x
        auto buf1bf16 = tmpbuf1bf16.Get<bfloat16_t>();
        auto buf2 = tmpbuf2.Get<float>();//ex
        auto buf3 = tmpbuf3.Get<float>(); //var
        auto buf4 = tmpbuf4.Get<float>();//eps
        auto buf5 = tmpbuf5.Get<float>();//
        float sumNN = 0;
        float sumval = 0;
        int32_t NN = H;
        int32_t h_groups = channels_per_group;
        int32_t len = H * W;
        float mean,var;
        eps = 1e-5;
        buf3(0) = sumNN;
        for(int32_t j = 0; j < h_groups ; j++)
        {
            DataCopyGM2UB<bfloat16_t>(buf1bf16, xGm[channels_per_group * H * W * progress + j * len], len);
            AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
            Cast(buf1, buf1bf16, AscendC::RoundMode::CAST_NONE, len);
            ReduceSum<float>(buf2, buf1, buf5, len);
            Add(buf3, buf2, buf3, 1);
            AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0);
        }
        sumNN = (float)buf3(0);
        mean = sumNN / (channels_per_group * H * W);
        buf5(0) = sumval;
        for(int32_t j = 0; j < h_groups; j++)
        { 
            DataCopyGM2UB<bfloat16_t>(buf1bf16, xGm[channels_per_group * H * W * progress + j * len], len);
            AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
            Cast(buf1, buf1bf16, AscendC::RoundMode::CAST_NONE, len);
            Duplicate<float>(buf2, mean, len);
            Sub(buf3, buf1, buf2, len);
            Mul(buf4, buf3, buf3, len);
            ReduceSum<float>(buf3, buf4, buf4, len);
            Add(buf5, buf3, buf5, 1);
            AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0);
        }
        sumval = (float)buf5(0);
        var = sumval / (channels_per_group * H * W);
        for(int32_t j = 0; j < h_groups; j++)
        {
            Duplicate<float>(buf3, var, len);
            AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
            DataCopyGM2UB<bfloat16_t>(buf1bf16, xGm[channels_per_group * H * W * progress + j * len], len);
            AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
            Cast(buf1, buf1bf16, AscendC::RoundMode::CAST_NONE, len);
            //x-ex
            Duplicate<float>(buf2, mean, len);
            Sub(buf1, buf1, buf2, len);
            Duplicate<float>(buf4, eps, len);
            Add(buf3, buf3, buf4, len);
            Sqrt(buf4, buf3, len);
            Div(buf1, buf1, buf4, len);
            Cast(buf1bf16, buf1, AscendC::RoundMode::CAST_RINT, len);
            AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
            DataCopyUB2GM<bfloat16_t>(yGm[channels_per_group * H * W * progress + j * len], buf1bf16, len);
            AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0);
        }

    }
    template <typename T>
    __aicore__ inline void DataCopyGM2UB(const LocalTensor<T> &dstTensor, const GlobalTensor<T> &srcTensor, const uint32_t &calCount)
    {
        DataCopyExtParams copyParams{1, calCount * static_cast<uint32_t>(sizeof(T)), 0, 0, 0};
        DataCopyPadExtParams<T> padParams{false, 0, 0, 0};

        DataCopyPad(dstTensor, srcTensor, copyParams, padParams);
    }

    template <typename T>
    __aicore__ inline void DataCopyUB2GM(const GlobalTensor<T> &dstTensor, const LocalTensor<T> &srcTensor, const uint32_t &calCount)
    {
        DataCopyExtParams copyParams{1, calCount * static_cast<uint32_t>(sizeof(T)), 0, 0, 0};
        DataCopyPad(dstTensor, srcTensor, copyParams);
    }
private:
    GlobalTensor<bfloat16_t> xGm;
    GlobalTensor<bfloat16_t> yGm;
    TPipe pipe;
    TBuf<QuePosition::VECCALC> tmpbuf1, tmpbuf1bf16, tmpbuf2, tmpbuf3, tmpbuf4, tmpbuf5;

    int32_t N, C, H, W, DT;
    int32_t num_groups, num_channels, channels_per_group;
    float eps;
    bool affine;
};

class GroupNormKernelhalf
{
public:
    __aicore__ inline GroupNormKernelhalf() {} 

    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, int32_t num_groups, int32_t num_channels, int32_t channels_per_group,
                                float eps, bool affine, int32_t N, int32_t C, int32_t H, int32_t W)
    {
        this->num_groups = num_groups;
        this->num_channels = num_channels;
        this->channels_per_group = channels_per_group;
        this->eps = eps;
        this->affine = affine;
        this->N = N;
        this->C = C;
        this->H = H;
        this->W = W;
        xGm.SetGlobalBuffer((__gm__ half *)x, N * C * H * W);
        yGm.SetGlobalBuffer((__gm__ half *)y, N * C * H * W);
        int32_t NN = 2;
        if (TILING_KEY_IS(7)||TILING_KEY_IS(8))
        {
            pipe.InitBuffer(tmpbuf1, this->channels_per_group * H * sizeof(float));
            pipe.InitBuffer(tmpbuf1fp16, this->channels_per_group * H * sizeof(half));
            pipe.InitBuffer(tmpbuf2, this->channels_per_group * H * sizeof(float));
            pipe.InitBuffer(tmpbuf3, this->channels_per_group * H * sizeof(float));
            pipe.InitBuffer(tmpbuf4, this->channels_per_group * H * sizeof(float));
            pipe.InitBuffer(tmpbuf5, this->channels_per_group * H * sizeof(float));
        }
        if (TILING_KEY_IS(9))
        {
            pipe.InitBuffer(tmpbuf1, NN * H * W * sizeof(float));
            pipe.InitBuffer(tmpbuf1fp16, NN * H * W * sizeof(half));
            pipe.InitBuffer(tmpbuf2, NN * H * W * sizeof(float));
            pipe.InitBuffer(tmpbuf3, NN * H * W * sizeof(float));
            pipe.InitBuffer(tmpbuf4, NN * H * W * sizeof(float));
            pipe.InitBuffer(tmpbuf5, NN * H * W * sizeof(float));
        }
    }
    __aicore__ inline void Process7()
    {
        for (int32_t i = 0; i < N * num_groups; i++)
        {
            compute5(i);
        }
    }
    __aicore__ inline void Process8()
    {
        for (int32_t i = 0; i < N * num_groups; i++)
        {
            compute5(i);
        }
    }
    __aicore__ inline void Process9()
    {  
        for (int32_t i = 0; i < N * num_groups; i++)
        {   
            compute6(i);
        }
    }
private:
    __aicore__ inline void compute5(int32_t progress)
    {
        auto buf1 = tmpbuf1.Get<float>();//x
        auto buf1fp16 = tmpbuf1fp16.Get<half>();
        auto buf2 = tmpbuf2.Get<float>();//ex
        auto buf3 = tmpbuf3.Get<float>(); //var
        auto buf4 = tmpbuf4.Get<float>();//eps
        auto buf5 = tmpbuf5.Get<float>();//eps
        int32_t len = channels_per_group * H;
        float mean,var,sum2;
        eps = 1e-5;
        DataCopyGM2UB<half>(buf1fp16, xGm[len * progress], len);
        AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
        AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
        Cast(buf1, buf1fp16, AscendC::RoundMode::CAST_NONE, len);
        ReduceSum<float>(buf2, buf1, buf5, len);
        buf3(0) = len;
        Div(buf3, buf2, buf3, 1);
        mean = buf3(0);

        Mul(buf2, buf1, buf1, len);
        ReduceSum<float>(buf4, buf2, buf5, len);
        sum2 = buf4(0);
        var = (sum2 - len * mean * mean) / len;
        Duplicate<float>(buf3, var, len);
        AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
        AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
        //x-ex
        Duplicate<float>(buf2, mean, len);
        Sub(buf1, buf1, buf2, len);
        //var+eps
        Duplicate<float>(buf4, eps, len);
        Add(buf3, buf3, buf4, len);
        Sqrt(buf4, buf3, len);
        Div(buf1, buf1, buf4, len);
        Cast(buf1fp16, buf1, AscendC::RoundMode::CAST_RINT, len);
        AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
        AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
        DataCopyUB2GM<half>(yGm[len * progress], buf1fp16, len);
        AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0);
        AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0);
    }
private:
    __aicore__ inline void compute6(int32_t progress)
    {
        auto buf1 = tmpbuf1.Get<float>();//x
        auto buf1fp16 = tmpbuf1fp16.Get<half>();
        auto buf2 = tmpbuf2.Get<float>();//ex
        auto buf3 = tmpbuf3.Get<float>(); //var
        auto buf4 = tmpbuf4.Get<float>();//epss
        auto buf5 = tmpbuf5.Get<float>();//
        float sumNN = 0;
        float sum2 = 0;
        int32_t NN = 2;
        int32_t T = channels_per_group * H * W;
        int32_t h_groups = channels_per_group / NN;
        int32_t len = NN * H * W;
        int32_t lentail = (channels_per_group % NN) * H * W;
        int32_t lenTemp;
        int32_t Times = 0;
        if(lentail == 0){
            Times = h_groups;
        }
        else{
            Times = h_groups + 1;
        }
        float mean,var;
        eps = 1e-5;
        buf3(0) = sumNN;
        buf4(0) = sum2;
        for(int32_t j = 0; j < Times ; j++)
            {
                if(j == h_groups){
                    lenTemp = lentail;
                }else{
                    lenTemp = len;
                }
                DataCopyGM2UB<half>(buf1fp16, xGm[T * progress + j * len], lenTemp);
                AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
                AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
                Cast(buf1, buf1fp16, AscendC::RoundMode::CAST_NONE, lenTemp);
                ReduceSum<float>(buf2, buf1, buf5, lenTemp);
                Add(buf3, buf2, buf3, 1);
                AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0);
                AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0);
                //jiade
                Mul(buf2, buf1, buf1, lenTemp);
                ReduceSum<float>(buf5, buf2, buf2, lenTemp);
                Add(buf4, buf5, buf4, 1);
            }
        sumNN = (float)buf3(0);
        sum2 = (float)buf4(0);
        mean = sumNN / (T);
        var = (sum2 - T * mean * mean) / T;
        // buf5(0) = sumval;
        // for(int32_t j = 0; j < Times; j++)
        // { 
        //     if(j == h_groups){
        //         lenTemp = lentail;
        //     }else{
        //         lenTemp = len;
        //     }
        //     DataCopyGM2UB<half>(buf1fp16, xGm[channels_per_group * H * W * progress + j * len], lenTemp);
        //     AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
        //     AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
        //     Cast(buf1, buf1fp16, AscendC::RoundMode::CAST_NONE, lenTemp);
        //     Duplicate<float>(buf2, mean, lenTemp);
        //     Sub(buf3, buf1, buf2, lenTemp);
        //     Mul(buf4, buf3, buf3, lenTemp);
        //     ReduceSum<float>(buf3, buf4, buf4, lenTemp);
        //     Add(buf5, buf3, buf5, 1);
        //     AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0);
        //     AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0);
        // }
        // sumval = (float)buf5(0);
        // var = sumval / (channels_per_group * H * W);
        for(int32_t j = 0; j < Times; j++)
        {
            if(j == h_groups){
                lenTemp = lentail;
            }else{
                lenTemp = len;
            }
            Duplicate<float>(buf3, var, lenTemp);
            AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
            DataCopyGM2UB<half>(buf1fp16, xGm[T * progress + j * len], lenTemp);
            AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
            Cast(buf1, buf1fp16, AscendC::RoundMode::CAST_NONE, lenTemp);
            //x-ex
            Duplicate<float>(buf2, mean, lenTemp);
            Sub(buf1, buf1, buf2, lenTemp);
            Duplicate<float>(buf4, eps, lenTemp);
            Add(buf3, buf3, buf4, lenTemp);
            Sqrt(buf4, buf3, lenTemp);
            Div(buf1, buf1, buf4, lenTemp);
            Cast(buf1fp16, buf1, AscendC::RoundMode::CAST_RINT, lenTemp);
            AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
            DataCopyUB2GM<half>(yGm[T * progress + j * len], buf1fp16, lenTemp);
            AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0);
            AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(EVENT_ID0);
        }
    }
    template <typename T>
    __aicore__ inline void DataCopyGM2UB(const LocalTensor<T> &dstTensor, const GlobalTensor<T> &srcTensor, const uint32_t &calCount)
    {
        DataCopyExtParams copyParams{1, calCount * static_cast<uint32_t>(sizeof(T)), 0, 0, 0};
        DataCopyPadExtParams<T> padParams{false, 0, 0, 0};

        DataCopyPad(dstTensor, srcTensor, copyParams, padParams);
    }

    template <typename T>
    __aicore__ inline void DataCopyUB2GM(const GlobalTensor<T> &dstTensor, const LocalTensor<T> &srcTensor, const uint32_t &calCount)
    {
        DataCopyExtParams copyParams{1, calCount * static_cast<uint32_t>(sizeof(T)), 0, 0, 0};
        DataCopyPad(dstTensor, srcTensor, copyParams);
    }
private:
    GlobalTensor<half> xGm;
    GlobalTensor<half> yGm;
    TPipe pipe;
    TBuf<QuePosition::VECCALC> tmpbuf1, tmpbuf1fp16, tmpbuf2, tmpbuf3, tmpbuf4, tmpbuf5;

    int32_t N, C, H, W, DT;
    int32_t num_groups, num_channels, channels_per_group;
    float eps;
    bool affine;
};
extern "C" __global__ __aicore__ void group_norm(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
{
    GET_TILING_DATA(tiling_data, tiling);
    if(TILING_KEY_IS(1)||TILING_KEY_IS(2)||TILING_KEY_IS(3)){
    // TODO: user kernel impl
        GroupNormKernelfloat op;
        op.Init(x, y, workspace, tiling_data.num_groups, tiling_data.num_channels, tiling_data.channels_per_group,
                tiling_data.eps, tiling_data.affine, tiling_data.N, tiling_data.C, tiling_data.H, tiling_data.W);
        if (TILING_KEY_IS(1))
        {
            op.Process1();
        }
        if (TILING_KEY_IS(2))
        {
            op.Process2();
        }
        if (TILING_KEY_IS(3))
        {
            op.Process3();
        }
    }
    if(TILING_KEY_IS(4)||TILING_KEY_IS(5)||TILING_KEY_IS(6)){
    // TODO: user kernel impl
        GroupNormKernelbf16 op;
        op.Init(x, y, workspace, tiling_data.num_groups, tiling_data.num_channels, tiling_data.channels_per_group,
                tiling_data.eps, tiling_data.affine, tiling_data.N, tiling_data.C, tiling_data.H, tiling_data.W);
        if (TILING_KEY_IS(4))
        {
            op.Process4();
        }
        if (TILING_KEY_IS(5))
        {
            op.Process5();
        }
        if (TILING_KEY_IS(6))
        {
            op.Process6();
        }
    }
    if(TILING_KEY_IS(7)||TILING_KEY_IS(8)||TILING_KEY_IS(9)){
    // TODO: user kernel impl
        GroupNormKernelhalf op;
        op.Init(x, y, workspace, tiling_data.num_groups, tiling_data.num_channels, tiling_data.channels_per_group,
                tiling_data.eps, tiling_data.affine, tiling_data.N, tiling_data.C, tiling_data.H, tiling_data.W);
        if (TILING_KEY_IS(7))
        {
            op.Process7();
        }
        if (TILING_KEY_IS(8))
        {
            op.Process8();
        }
        if (TILING_KEY_IS(9))
        {
            op.Process9();
        }
    }
}